#include "cuda.h"

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>

#include <fstream>
#include <time.h>
#include <iomanip>
#include <queue>
#include <assert.h>
#include <stdlib.h>
using namespace std;

__global__ void initKernel(devNode** dev_ptr, int* matrix_d, int cycle, int depth){
    
    int tid = blockIdx.x * 512 + threadIdx.x;
    if(tid < depth){
        if(matrix_d[cycle*depth+tid] == 0){
            dev_ptr[0][tid].setLogic(ZERO);
        }else{
            dev_ptr[0][tid].setLogic(ONE);
        }
    }
}

__global__ void simKernel(devNode** dev_ptr, int level, int depth){
    
    int tid = blockIdx.x * 512 + threadIdx.x;
    int logic = X;
    int gate;
    int fLevelA, fLevelB, fLevelC, fLevelD; 
    int faninA, faninB, faninC, faninD;
    
    if(tid < depth){
        gate = dev_ptr[level][tid].getGate();
        if(gate == INVX1 || gate == INVX2 || gate == INVX4 || gate == INVX8){
            fLevelA = dev_ptr[level][tid].getLevel(0);
            faninA = dev_ptr[level][tid].getFanin(0);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic());
            
        }else if(gate == AND2X1 || gate == AND2X2 || gate == HAC){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = dev_ptr[fLevelA][faninA].getLogic() & dev_ptr[fLevelB][faninB].getLogic();
            
        }else if(gate == NAND2X1){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic() & dev_ptr[fLevelB][faninB].getLogic());
            
        }else if(gate == NAND3X1){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);  
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic() & dev_ptr[fLevelB][faninB].getLogic() &
                    dev_ptr[fLevelC][faninC].getLogic());
            
        }else if(gate == OR2X1 || gate == OR2X2){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = dev_ptr[fLevelA][faninA].getLogic() | dev_ptr[fLevelB][faninB].getLogic();
            
        }else if(gate == NOR2X1){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic() | dev_ptr[fLevelB][faninB].getLogic());
            
        }else if(gate == NOR3X1){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);  
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic() | dev_ptr[fLevelB][faninB].getLogic() |
                    dev_ptr[fLevelC][faninC].getLogic());
            
        }else if(gate == XOR2X1 || gate == HAS){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = dev_ptr[fLevelA][faninA].getLogic() ^ dev_ptr[fLevelB][faninB].getLogic();
            
        }else if(gate == XNOR2X1){
            fLevelB = dev_ptr[level][tid].getLevel(0);
            faninB = dev_ptr[level][tid].getFanin(0);
            fLevelA = dev_ptr[level][tid].getLevel(1);
            faninA = dev_ptr[level][tid].getFanin(1);
            logic = ~(dev_ptr[fLevelA][faninA].getLogic() ^ dev_ptr[fLevelB][faninB].getLogic());
            
        }else if(gate == AOI21X1){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);  
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = ~((dev_ptr[fLevelA][faninA].getLogic() & dev_ptr[fLevelB][faninB].getLogic()) |
                    dev_ptr[fLevelC][faninC].getLogic());
            
        }else if(gate == AOI22X1){
            fLevelD = dev_ptr[level][tid].getLevel(0);
            faninD = dev_ptr[level][tid].getFanin(0);
            fLevelC = dev_ptr[level][tid].getLevel(1);
            faninC = dev_ptr[level][tid].getFanin(1);  
            fLevelB = dev_ptr[level][tid].getLevel(2);
            faninB = dev_ptr[level][tid].getFanin(2);
            fLevelA = dev_ptr[level][tid].getLevel(3);
            faninA = dev_ptr[level][tid].getFanin(3);
            logic = ~((dev_ptr[fLevelA][faninA].getLogic() & dev_ptr[fLevelB][faninB].getLogic()) |
                    (dev_ptr[fLevelC][faninC].getLogic() & dev_ptr[fLevelD][faninD].getLogic()));
            
        }else if(gate == OAI21X1){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = ~((dev_ptr[fLevelA][faninA].getLogic() | dev_ptr[fLevelB][faninB].getLogic()) &
                    dev_ptr[fLevelC][faninC].getLogic());
            
        }else if(gate == OAI22X1){
            fLevelD = dev_ptr[level][tid].getLevel(0);
            faninD = dev_ptr[level][tid].getFanin(0);
            fLevelC = dev_ptr[level][tid].getLevel(1);
            faninC = dev_ptr[level][tid].getFanin(1);  
            fLevelB = dev_ptr[level][tid].getLevel(2);
            faninB = dev_ptr[level][tid].getFanin(2);
            fLevelA = dev_ptr[level][tid].getLevel(3);
            faninA = dev_ptr[level][tid].getFanin(3);
            logic = ~((dev_ptr[fLevelA][faninA].getLogic() | dev_ptr[fLevelB][faninB].getLogic()) &
                    (dev_ptr[fLevelC][faninC].getLogic() | dev_ptr[fLevelD][faninD].getLogic()));
        }else if(gate == BUFX2 || gate == BUFX4){
            fLevelA = dev_ptr[level][tid].getLevel(0);
            faninA = dev_ptr[level][tid].getFanin(0);
            logic = dev_ptr[fLevelA][faninA].getLogic();
        }else if(gate == FAS){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = (dev_ptr[fLevelA][faninA].getLogic()) ^ (dev_ptr[fLevelB][faninB].getLogic()) ^
                    (dev_ptr[fLevelC][faninC].getLogic());
        }else if(gate == FAC){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            logic = ((dev_ptr[fLevelA][faninA].getLogic()) & (dev_ptr[fLevelB][faninB].getLogic())) |
                    ((dev_ptr[fLevelA][faninA].getLogic()) & (dev_ptr[fLevelC][faninC].getLogic())) |
                    ((dev_ptr[fLevelB][faninB].getLogic()) & (dev_ptr[fLevelC][faninC].getLogic()));
        }else if(gate == MUX2X1){
            fLevelC = dev_ptr[level][tid].getLevel(0);
            faninC = dev_ptr[level][tid].getFanin(0);
            fLevelB = dev_ptr[level][tid].getLevel(1);
            faninB = dev_ptr[level][tid].getFanin(1);
            fLevelA = dev_ptr[level][tid].getLevel(2);
            faninA = dev_ptr[level][tid].getFanin(2);
            if(dev_ptr[fLevelC][faninC].getLogic() == ZERO){
                logic = dev_ptr[fLevelB][faninB].getLogic();
            }else{
                logic = dev_ptr[fLevelA][faninA].getLogic();
            }
        }
        
        dev_ptr[level][tid].setLogic((LType)(logic & 3));
    }
}


__global__ void ffKernel(devNode** dev_ptr, int depth){
    
    int tid = blockIdx.x * 512 + threadIdx.x;
    int level, offset;
    if(tid < depth){
        if(dev_ptr[0][tid].getType() == PPI){
            level = dev_ptr[0][tid].getLevel(0);
            offset = dev_ptr[0][tid].getFanin(0);
            dev_ptr[0][tid].setLogic(dev_ptr[level][offset].getLogic());
        }
    }
}

//check if node is ready to compute its offset
int nodeOffsetReady(Node* np){
    int flag = 1;
    vector<Node*>::iterator it;
    for(it = np->upperNodes.begin(); it < np->upperNodes.end(); it++){
        if((*it)->offset == -1){
            flag = 0;
            break;
        }
    }
    return flag;
}

void transform(){
    
    queue<Node*> nodeQ;
    vector<Node*>::iterator it, itf;
    map<Node*, Node*>::iterator itt;
    Node* np;
    devNode* dp;
    
    module* mp = moduleAll.back();
    Node** npp = mp->nInput;
    
    gMatrix = new thrust::host_vector<devNode> [maxLevel + 1];
	ptrMatrix = new vector<Node*> [maxLevel + 1];
    
    //process all input nodes of top module
    for(int i = 0; i < mp->inputNum; i++){
		ptrMatrix[0].push_back(npp[i]);
        npp[i]->offset = gMatrix[0].size();
		gMatrix[0].push_back(devNode(npp[i]->modIndex, npp[i]->index, npp[i]->type, npp[i]->gate, npp[i]->fanin));
        for(it = npp[i]->downNodes.begin(); it < npp[i]->downNodes.end(); it++){
            if(nodeOffsetReady((*it))){
                nodeQ.push((*it));
            }
        }
    }
    
    //process flip flop of all modules
    for(int i = 0; i < moduleAll.size(); i++){
        mp = moduleAll[i];
        for(itt = mp->ffMap.begin(); itt != mp->ffMap.end(); itt++){
            np = (*itt).second;
			ptrMatrix[0].push_back(np);
            np->offset = gMatrix[0].size();
			dp = new devNode(np->modIndex, np->index, np->type, np->gate, np->fanin);
            gMatrix[0].push_back(*dp);
            for(it = np->downNodes.begin(); it < np->downNodes.end(); it++){
                if(nodeOffsetReady((*it))){
                    nodeQ.push((*it));
                }
            }
        }
    }
    
    //process dummy nodes of all modules
    for(int i = 0; i < moduleAll.size(); i++){
        mp = moduleAll[i];
        for(itf = mp->nFloat.begin(); itf < mp->nFloat.end(); itf++){
			ptrMatrix[0].push_back((*itf));
            (*itf)->offset = gMatrix[0].size();
			gMatrix[0].push_back(devNode((*itf)->modIndex, (*itf)->index, (*itf)->type, (*itf)->gate, (*itf)->fanin, (*itf)->logic));
            for(it = (*itf)->downNodes.begin(); it < (*itf)->downNodes.end(); it++){
                if(nodeOffsetReady((*it))){
                    nodeQ.push((*it));
                }
            }
        }
    }
    
    //process the queue
    while (nodeQ.size() != 0) {
        np = nodeQ.front();
        nodeQ.pop();
        np->offset = gMatrix[np->level].size();
		dp = new devNode(np->modIndex, np->index, np->type, np->gate, np->fanin);
        for(int i = 0; i < np->fanin; i++){
            dp->setLevel(i, np->upperNodes[i]->level);
            dp->setFanin(i, np->upperNodes[i]->offset);
        }
        gMatrix[np->level].push_back(*dp);
		ptrMatrix[np->level].push_back(np);
        for(it = np->downNodes.begin(); it < np->downNodes.end(); it++){
            if(nodeOffsetReady((*it))){
                nodeQ.push((*it));
            }
        }
    }
    
    //process flip flop's fanin of all modules
    for(int i = 0; i < moduleAll.size(); i++){
        mp = moduleAll[i];
        for(itt = mp->ffMap.begin(); itt != mp->ffMap.end(); itt++){
            np = (*itt).second;
            gMatrix[np->level][np->offset].setLevel(0, (*itt).first->level);
            gMatrix[np->level][np->offset].setFanin(0, (*itt).first->offset);
        }
    }
}

int findBias(int level, int index){
    
    int bias = -1;
    for(int i = 0; i < gMatrix[level].size(); i++){
        if(gMatrix[level][i].getIndex() == index){
            bias = i;
            break;
        }
    }
    return bias;
}

//sequential update fanin
void updateFanin(){
    
    Node* np;
    int bias = -1;
    int level;
    
    for(int i = 1; i <= maxLevel; i++){
        for(int j = 0; j < gMatrix[i].size(); j++){
            np = nodeAll[gMatrix[i][j].getIndex()];
            for(int k = 0; k < np->fanin; k++){
                level = i - 1;
                while((bias = findBias(level, np->upperNodes[k]->index)) == -1
                        && level >= 0){
                    level--;
                }
                gMatrix[i][j].setFanin(k, bias);
                gMatrix[i][j].setLevel(k, level);
            }
        }
    }
}

void cudaLogicSim(){
    
    ofstream result;
    int depth, grid;
    devNode** temp;
    int* matrix_dev;
    module* mp = moduleAll.back();
    test = new thrust::host_vector<devNode> [maxLevel+1];
    
    cudaMalloc((void**)&matrix_dev, (mp->inputNum)*MAXCYCLE*sizeof(int));
    cudaMemcpy(matrix_dev, matrix, (mp->inputNum)*MAXCYCLE*sizeof(int), cudaMemcpyHostToDevice);
    
    for(int i = 0; i < gMatrix[0].size(); i++){
        if(gMatrix[0][i].getType() == PPI){
            gMatrix[0][i].setLogic(ZERO);
        }
    }
    
    //copy all vectors in gMatrix from main memory into device memory
    stem_host.resize(maxLevel + 1);
    for(int i = 0; i <= maxLevel; i++){
        cudaMalloc((void**)&stem_host[i], gMatrix[i].size() * sizeof(devNode));
        thrust::device_ptr<devNode> dev_ptr(stem_host[i]); //change raw pointer into thrust pointer
        thrust::copy(gMatrix[i].begin(), gMatrix[i].end(), dev_ptr); //thrust copy is used
    }
    
    //copy reference of all vectors from main memory into device memory
    cudaMalloc(&temp, stem_host.size()*sizeof(devNode*));
    devNode** raw_ptr = thrust::raw_pointer_cast(&stem_host[0]); //change thrust pointer into raw pointer
    cudaMemcpy(temp, raw_ptr, stem_host.size()*sizeof(devNode*), cudaMemcpyHostToDevice);
    
    //issue series of kernel calls
    for(int i = 0; i < MAXCYCLE; i++){
//        cout<<"cuda simulate cycle "<<i<<endl;
        depth = mp->inputNum;
        grid = depth/512 + 1;
        initKernel<<<grid, 512>>>(temp, matrix_dev, i, depth);
        for(int j = 1; j <= maxLevel; j++){
            depth = gMatrix[j].size();
            grid = depth/512 + 1;
            simKernel<<<grid, 512>>>(temp, j, depth);
        }
        depth = gMatrix[0].size();
        grid = depth/512 + 1;
        if(i != MAXCYCLE-1){
            ffKernel<<<grid, 512>>>(temp, depth);
        }
    }
    
    //copy all results from device memory back to main memory and then print
    result.open("cuda.txt");
    for(int i = 0; i <= maxLevel; i++){
        thrust::device_ptr<devNode> dev_ptr_new(stem_host[i]);
        test[i].resize(gMatrix[i].size());
//        thrust::copy_n(dev_ptr_new, gMatrix[i].size(), test[i].begin()); //thrust copy is used
    }
    
    for(int i = 0; i <= maxLevel; i++){
        for(int j = 0; j < gMatrix[i].size(); j++){
            result<<test[i][j].getIndex()<<"->"<<lMap[test[i][j].getLogic()]<<" ";
        }
        result<<endl;
    }
    result.close();
    
    //free all memory allocated on device memory
    for(int i = 0; i <= maxLevel; i++){
        cudaFree(stem_host[i]);
    }
    cudaFree(temp);
    cudaFree(matrix_dev);
}

void cudaVerify(){
    
    Node* np;
    int flag = 1;
    
    cout<<"cuda verify is starting ..."<<endl;
    for(int i = 0; i < moduleAll.size(); i++){
        for(int j = 0; j < moduleAll[i]->allNum; j++){
            np = moduleAll[i]->nAll[j];
//            cout<<np->name<<endl;
            if(np->level != -1) np->cudaLogic = test[np->level][np->offset].getLogic();
            if(np->cudaLogic != np->logic) flag = 0;
        }
    }
    
    if(flag){
        cout<<"Pass!"<<endl;
    }else{
        cout<<"Fail!"<<endl;
    }
}

void cudaSimTest(int a, int b){
    
    long long result = 0;
    int temp;
    int tempa = a;
    int tempb = b;
    time_t start, end;
    double diff;
    
    for(int i = 0; i < 24; i++){
        temp = a & 1;
        if(temp == 0){
            gMatrix[0][i].setLogic(ZERO);
        }else{
            gMatrix[0][i].setLogic(ONE);
        }
        a = a >> 1;
    }
    
    for(int i = 24; i < 48; i++){
        temp = b & 1;
        if(temp == 0){
            gMatrix[0][i].setLogic(ZERO);
        }else{
            gMatrix[0][i].setLogic(ONE);
        }
        b = b >> 1;
    }
    
    time(&start);
    cudaLogicSim();
    time(&end);
    diff = difftime(end, start);
    
    for(int i = 47; i >= 0; i--){
        if(nodeOutput[i]->cudaLogic == ZERO){
            result = result & 0xfffffffffffffffeLL;
        }else if(nodeOutput[i]->cudaLogic == ONE){
            result = result | 1;
        }
        result = result << 1;
    }
    
    cout<<tempa<<" * "<<tempb<<" = "<<result<<endl;
    printf("GPU took %.5f seconds to do logic simulation\n", diff);
    
}
